//这个程序我们打算在内核函数中记录时间戳，然后在CPU逻辑中查看第一个线程和最后一个线程运行的时间。
#include<stdio.h>
#include<stdlib.h>
#include<string.h>

//就传入两个数组，一个数组记录每一个block开始的时间，一个数组记录每一个block结束的时间
//使用clock函数来进行计时，这个函数的返回值不是时间，而是每个线程在执行代码用了多少个时钟周期。
//所以时钟周期还要除以GPU频率才能算出真正的时间
//https://blog.csdn.net/adi_1987/article/details/53465370
//为了方式编译器的优化，我们才里面申请一段共享内存，并且将共享内存的数据求和，并且将结果放到一个数组中
__global__  void gpu_function(clock_t* start, clock_t* stop, int* input, int* result_arr){
    //每个线程块的第一个线程的记录线程
    if(threadIdx.x == 0){
        start[blockIdx.x] = clock();
    }

    //利用共享内存将在全局内存的输入数据取入
    __shared__ int temp[192];

    temp[threadIdx.x] = input[threadIdx.x];

    __syncthreads();
    
    //这里执行计算
    int i = 0;
    int result = 0;

    for(i = 0; i < 192; i++){
        result = result + temp[i];
    }

    //将数据放到全局内存中
    if(threadIdx.x == 2){
        result_arr[blockIdx.x] = result;
    }

    //这里执行同步
    __syncthreads();

    //这里记录结束的时间戳
    if(threadIdx.x == 0){
        stop[blockIdx.x] = clock();
    }
}

//我们的显卡一共有15个SM，每个SM有192个sp，我们打算运行30个线程块，每个线程块开启192个线程。
int main(){
    //这里激活内核函数
    dim3 grid_dimention(30);
    dim3 block_dimention(192);

    clock_t* start = NULL;
    clock_t* stop = NULL;

    cudaMallocHost((void **)&start, 30*sizeof(clock_t));
    cudaMallocHost((void **)&stop, 30*sizeof(clock_t));

    //这里申请两个记录时间戳的数组来
    clock_t* device_start = NULL;
    clock_t* device_stop = NULL;

    cudaMalloc((void **)&device_start, 30*sizeof(clock_t));
    cudaMalloc((void **)&device_start, 30*sizeof(clock_t));

    //输入数组指针
    int* input;
    //这里申请一个数组，这个数组里面有0-191个数字
    cudaMallocHost((void **)&input, 192 * sizeof(int));
    
    int i;

    for(i = 0; i < 192; i++){
        input[i] = i;
    }

    cudaEvent_t start_event;
    cudaEvent_t stop_event;
    cudaEventCreate(&start_event);
    cudaEventCreate(&stop_event);

    cudaEventRecord(start_event, 0);

    //这里申请输出空间
    int *device_result;
    cudaMalloc((void **)&device_result, 30 * sizeof(int));
    int *result;
    cudaMallocHost((void **)&result, 30 * sizeof(int));

    //将输入数组拷贝到显卡上
    int* device_input;

    cudaMalloc((void **)&device_input, 192 * sizeof(int));
    cudaMemcpy(device_input, input, 192 * sizeof(int), cudaMemcpyHostToDevice);

    

    gpu_function<<<grid_dimention, block_dimention>>>(device_start, device_start, device_input, device_result);


    cudaThreadSynchronize();
    
    //这里将数据拷出
    cudaMemcpy(start, device_start, 30*sizeof(clock_t), cudaMemcpyDeviceToHost);
    cudaMemcpy(stop, device_stop, 30*sizeof(clock_t), cudaMemcpyDeviceToHost);
    cudaMemcpy(result, device_result, 30*sizeof(int), cudaMemcpyDeviceToHost);

    cudaEventRecord(stop_event, 0);

    cudaEventSynchronize(stop_event);

    //记录时间
    float time;
    cudaEventElapsedTime(&time, start_event, stop_event);

    printf("基于事件的计时方式：%f\n", time);


    //这里获取GPU的主频，使用的方法就是cudaDeviceProp
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    
    int gpu_rate = prop.clockRate;

    //查找最早的开始时间和最后结束的时间
    clock_t min_clock = start[0];
    clock_t max_clock = stop[0];

    //选出第一个block的开始时间和最后一个block的结束时间
    for(i = 1; i < 30; i++){
        if(min_clock < start[i]){
            min_clock = start[i];
        }

        if(max_clock > stop[i]){
            max_clock = stop[i];
        }
    }

    printf("%ld\n", (max_clock-min_clock)/gpu_rate);

    cudaFreeHost(start);
    cudaFreeHost(stop);
    cudaFreeHost(result);
    cudaFreeHost(input);

    cudaEventDestroy(start_event);
    cudaEventDestroy(stop_event);

    cudaFree(device_start);
    cudaFree(device_stop);
    cudaFree(device_input);
    cudaFree(device_result);

    return 0;
}